Make internal host_allocator stream ordered - #23574

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

/ok to test 8ebdb6c

Summary

This PR attempted to replace the cudaMallocHost-backed pinned pool with a CUDA native host memory pool
(cudaMemPoolCreate with cudaMemAllocationTypePinned + cudaMemLocationTypeHost, allocating via cudaMallocFromPoolAsync
/ freeing via cudaFreeAsync). The goal was twofold:

1. Remove stream.synchronize() from rmm_host_allocator::allocate() — the CUDA pool guarantees it only recycles blocks
   with no outstanding GPU accesses, so no external sync is needed before CPU writes.
2. Stream-ordered deallocation — cudaFreeAsync defers the reclaim until the stream reaches the free point, meaning
   callers can destroy a host_vector immediately after enqueuing a device copy without needing to sync first.

The implementation compiled and all local tests passed on CUDA 12.8 / sm_86 (Ampere). However, CI exposed three
failure modes, all sharing the same root cause:

Root cause

cudaMallocFromPoolAsync with cudaMemLocationTypeHost does not behave as a drop-in replacement for cudaMallocHost.
cudaMallocHost registers pinned pages through CUDA's unified addressing system, giving them proper pointer attributes
(cudaMemoryTypeHost + a valid devicePointer) that all CUDA memory transfer paths understand. The native host pool
allocates page-locked memory and can be made accessible to the GPU via cudaMemPoolSetAccess, but that grants peer
access, not the same unified-addressing registration. As a result:

• cudaMemcpyAsync (H→D / D→H DMA) fails on some GPU/driver combinations because the DMA engine doesn't find the
  expected pinned registration.
• cudaMemcpyDeviceToDevice fails unconditionally when one pointer is pool host memory, because D→D DMA expects both
  endpoints in device DRAM — the peer access mapping is invisible to that path.

Specific failures

Failure: cudaMemPoolCreate
Environment: CUDA 12.2.2 amd64
Error: cudaErrorInvalidValue — host memory pools require a minimum driver version not present on that runner
────────────────────────────────────────
Failure: cudaMemcpyAsync in IO paths (HybridScanFiltersTest, ParquetTest.ChunkedOperations)
Environment: CUDA 12.9+ amd64
Error: cudaErrorIllegalAddress — DMA path rejects pool host memory on those GPUs
────────────────────────────────────────
Failure: TDigestMergeTest (groupby + reductions)
Environment: Local CUDA 12.8
Error: cudaErrorInvalidValue in thrust::copy D→D — tdigest merge uses cudf::get_pinned_memory_resource() to allocate
device_uvectors in pinned memory and copies into them with cudaMemcpyDeviceToDevice, which rejects pool host pointers

Future path

There is a device attribute — cudaDevAttrHostMemoryPoolsSupported (144, added in CUDA 13.x headers) — that correctly
gates whether this feature is usable. The right implementation would query this attribute at pool construction time
and fall back to the cudaMallocHost-backed pool when it returns 0. Even then, the cudaMemcpyDeviceToDevice
incompatibility would need to be addressed separately, either by using cudaMemcpyDefault everywhere, or by ensuring
pool host memory goes through a different allocation path than the one device_uvector uses for its unified-addressing
trick. A safe minimum would likely be CUDA 13.x + sm_80+ (Ampere) + cudaDevAttrHostMemoryPoolsSupported == 1.

We do need to replace cudaMemcpyDeviceToDevice with cudaMemcpyDefault. Where is that source code? Thrust or cuDF?

For CUDA 12.x we need to use a different set of APIs for pinned memory, I think. There is a bit of nuance here. See rapidsai/rmm#2054.

I opened #23605 to use cudaMemcpyDefault everywhere.

/ok to test